Skip to content

feat(rocm): grouped + non-grouped keep-quant expert GEMM — the MoE-path terminus (issue #41) - #523

Draft
VikashLoomba wants to merge 2 commits into
mudler:mainfrom
VikashLoomba:row/ROCM-GG-ON-MOE
Draft

feat(rocm): grouped + non-grouped keep-quant expert GEMM — the MoE-path terminus (issue #41)#523
VikashLoomba wants to merge 2 commits into
mudler:mainfrom
VikashLoomba:row/ROCM-GG-ON-MOE

Conversation

@VikashLoomba

Copy link
Copy Markdown
Contributor

Row

BACKEND-ROCM — the last kernel blocker for MoE-bearing models on discrete ROCm. Issue #41. Stacked on #509 (MoE combine/gate chain).

What changed

NEW src/vt/rocm/rocm_grouped_gemm.hip — ports the keep-quant expert GEMM family from cuda_quant_dot.cu 1:1:

  • kMatmulBTQuant (op 74, non-grouped) + kMatmulBTQuantGrouped (op 75, grouped over expert_ids)
  • Q8_0 / Q4_K / Q5_K / Q6_K weight formats (the ones the target GDN-MoE GGUFs carry); Q8_0 + Q8_K activation quantizers
  • DotQ8_0/DotQ4K/DotQ5K/DotQ6K superblocks, __dp4a → portable Dp4a (bit-identical integer core), __shfl_down_sync reduction
  • Registering the non-grouped op is what flips GgufQuantComputeAvailable() on ROCm — the grouped op alone leaves the loader dequantizing experts to bf16.

Cross-device case: all four formats vs the CPU keep-quant oracle, valid random blocks + real activations.

Evidence (4× gfx1100, ROCm 7.14, Release)

  • grouped-quant case: 16/16 assertions across Q8_0/Q4_K/Q5_K/Q6_K, NMSE ≤ 5e-4 vs the CPU keep-quant reference
  • ctest -R 'rocm|cross_device': 4/4; full ctest: only the 5 pre-existing host/lane failures (zero new)
  • E2E: Qwen3.6-35B-A3B Q4_K_M (21GB GDN-MoE GGUF) runs end to end on one gfx1100 with keep-quant active — op 74 AND op 75 resolve vt-native, zero CPU-ref fallback, correct output. Requires --max-num-seqs 1 to fit one 24GB card (the GDN state pool otherwise pushes past — a residency note, not a kernel defect; confirmed via a backend-Alloc instrumentation run showing genuine cumulative ~23.8 GiB).
  • preflight --staged + trailers green

Speed claims

  • This PR makes NO speed claim.

Honest gaps

  • The K-quant formats Q2_K/Q3_K and the IQ2/IQ3 family are NOT ported (throw loudly); the target models don't use them. They follow the same skeleton.
  • The per-call hipMalloc/hipFree activation scratch is correctness-grade; a queue-owned grow-only pool is a perf lever (decode-step churn).
  • The 35B fits only with --max-num-seqs 1 on 24GB; multi-GPU expert sharding or host streaming is the follow-on for bigger MoE / longer context.

…mbineGate) (mudler#41)

The next links in the generic MoE path after the router/silu-mul. Hand-
translated from cuda_moe.cu (MoeCombineKernel :473, MoeCombineGateKernel :555)
and the SharedExpertGate CPU oracle (cpu_ops.cpp:2387). Grid-stride, f32 math,
bf16/f32 dtype arms via boundary conversions; the combine-gate folds the
shared-expert sigmoid gate rounded through bf16 exactly as the donor.

Evidence (4x gfx1100, ROCm 7.14, Release):
- new MoE combine/gate cross-device case: 9/9 assertions (MoeCombineGate's
  oracle is the host-computed composite — no CPU op registration exists)
- ctest -R 'rocm|cross_device': 4/4
- full ctest: pre-existing failure set shrinks 7 -> 5; test_bench and
  test_capi now PASS (they failed at op 77 / the router dtype before the
  chain). test_loaded_engine_dense now fails only on the async-scheduling
  assertion (a lane capability gap, not a kernel throw).
- Named remaining blocker: the grouped quant expert GEMM
  (kMatmulBTQuantGrouped), the DeepSeek-V4 keep-quant family.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: pi:kimi-k3 [pi]
…ant, kMatmulBTQuantGrouped) — the MoE-path terminus (mudler#41)

The last kernel blocker for GDN-MoE models on discrete ROCm. Ports the
cuda_quant_dot.cu grouped + plain keep-quant GEMMs (QuantDotGemmGroupedKernel
:746 / QuantDotGemmKernel :706 / the Q8_0 + Q8_K activation quantizers) and the
DotQ8_0 / DotQ4K / DotQ5K / DotQ6K superblocks 1:1. Registers kMatmulBTQuant
(op 74) + kMatmulBTQuantGrouped (op 75) — registering the non-grouped op flips
GgufQuantComputeAvailable() so the keep-quant path actually activates on ROCm.
Integer dot cores via a portable Dp4a (bit-identical to __dp4a); the HW dot
instruction is a named perf lever.

Evidence (4x gfx1100, ROCm 7.14, Release):
- cross-device grouped-quant case: 16/16 (Q8_0/Q4_K/Q5_K/Q6_K, valid random
  blocks + real f32 activations, NMSE <= 5e-4 vs the CPU keep-quant oracle)
- ctest -R 'rocm|cross_device': 4/4; full ctest: the 5 pre-existing
  host/lane failures only (test_bench + test_capi pass since the MoE chain).
- E2E: Qwen3.6-35B-A3B Q4_K_M (21GB GDN-MoE GGUF) runs end to end on one
  gfx1100 with keep-quant active — ops 74 AND 75 resolve vt-native, zero
  CPU-ref fallback, correct output. Needs --max-num-seqs 1 to fit one 24GB
  card (the GDN state pool otherwise pushes past; residency note, not a
  kernel defect). A backend-Alloc instrumentation run confirmed the OOM was
  cumulative ~23.8 GiB = genuine capacity, not a kernel bug.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: pi:kimi-k3 [pi]
@localai-bot

Copy link
Copy Markdown
Collaborator

Reviewed as part of a sweep over the open external PRs. The GEMM port is the strongest part of this — I diffed DotQ4K/DotQ5K/DotQ6K against the donor line by line and DotQ6K is byte-for-byte; all four cited anchors in cuda_quant_dot.cu resolve. Registration goes through the existing seam.

One blocker, and it is in how the op is switched on rather than in the kernel.

Registering kMatmulBTQuant flips keep-quant loader-wide for formats this kernel cannot execute.

GgufQuantComputeAvailable() (gguf_keep_quant.cpp:74) is a boolean: OpRegistered(kMatmulBTQuant, current_device). FromEnv() sets p.keep_quant from it, and KeepQuantDType admits anything vt::cpu::HasQuantDotKernel accepts — which per cpu_quant_traits.cpp:30-98 is Q4_0, Q8_0, Q2_K, Q3_K, Q4_K, Q5_K, Q6_K, IQ2_XXS, IQ3_XXS, IQ2_S and MXFP4. Twelve formats.

rocm_grouped_gemm.hip:519,592 implements four and throws for the rest. CUDA survives the same boolean because cuda_quant_dot.cu:1841-1846 falls back to the CPU kernel on an unsupported dtype. This PR's own docs/USAGE.md edit states the opposite for ROCm: "On a discrete card there is no CPU fallback tier."

So a Q4_0 or Q2_K or IQ2 GGUF that loads and generates correctly on ROCm today would now keep its blocks quantized at load and throw vt rocm: matmul_bt_quant: unsupported weight dtype on the first forward pass — after the whole model is resident on the card. That is a functional regression of a working path, discovered at the worst possible moment, and it is not mentioned in the PR body.

Same flip has collateral: expand_nkkeep_f16 turns ON, so an F16 file weight stays kF16 and MatmulBTKernelRocm (rocm_matmul_hipblaslt.hip:446-450) accepts only bf16/bf16 or f32/f32 → vt rocm: matmul_bt: unsupported.

The fix is a per-dtype availability query rather than a boolean, or a refusal at load time rather than forward time. Related: the refusal message at :592-594 lists Q5_K on both the ported and the not-yet sides, and does not name Q4_0 or MXFP4 — the two formats that actually drive this.

Second thing, worth knowing before you invest further. The per-call hipMalloc/hipFree + hipStreamSynchronize (:478,488-489,500,515-516,545,558-559,572,587-588) is not only the decode-step churn you name honestly in the body. Both are illegal under hipGraph stream capture, which hard-blocks the ROCm decode-graph work in #473/#332 — the row worth ~3x decode. The donor uses EnsureScratch(), a per-stream grow-only pool, and never synchronizes; rocm_device_bind.h:21 shows the codebase already cares about capture-safety. (Also: qact leaks if Check() throws between the malloc and the free.)

Smaller: Dp4a at :168-172 unrolls into four scalar MACs where HIP exposes __dp4a natively (lowers to v_dot4_i32_i8 on gfx9/gfx10+) — that is the innermost loop of every expert GEMM, so roughly 4x the instruction count on the hottest path for a one-line change. And there is no spec: grep -rn 'ROCM-GG-ON-MOE' .agents/ docs/ returns nothing, so the scope, gates and stop conditions for this work are not written down anywhere, and the unported arms are not recorded as owed.

Note also that the non-grouped op — the one that carries this PR's headline mechanism — has no test at all; only kMatmulBTQuantGrouped is exercised. And both new cases guard with if (!OpAvailable(op, dt)) continue;, so deleting the RegisterOp line makes them skip silently and stay green, which means nothing currently proves the registration.

On CI: the four red checks are all infrastructure, none yours. pr-size and agent-record both die on base must be an ancestor of head (unrebased fork) — and that is now fixed on our side in #619, so a rebase should clear them. Windows is the known-broken arm.

No AMD hardware here, so your 16/16, the ctest results and the 35B e2e run could not be reproduced and I am not disputing them; the finding above is read from the loader chain end to end.

localai-bot pushed a commit that referenced this pull request Aug 13, 2026
…(W1, #332) (#473)

Implements the vt::Backend graph-capture seam on hipGraph (W1 of #332), mirroring
src/vt/cuda/cuda_backend.cu call for call.

Merged with the row's performance rationale REFUTED and recorded as such. W3
measured capture at +3.2% / +0.6% / -1.0%, not the ~2.2-3x §1 predicted, and the
spec's D7 plus the inline note in §1 now say so rather than leaving a live
rationale for the next agent to re-derive. That refutation was reported by the
contributor against their own interest, which is the behaviour this protocol
exists to produce.

Merged anyway on the seam argument, which is independent of the decode number:
graph capture had exactly one real implementation (CUDA), and a one-implementation
abstraction is unproven. hipGraph is the cheapest available second. Runtime cost
today is zero -- RocmPlatform does not override support_static_graph_mode(), so
nothing in the engine reaches the new code.

Review verified isolation three ways: every decode-graph call site ANDs
SupportsGraphCapture() with support_static_graph_mode(); rocm_backend.hip appears
zero times in a CPU build's compile_commands.json; and there is zero drift on
every touched file across the 194-commit gap. A mutation of the EndCaptureGraph
seam signature turns the new test red, so it genuinely guards seam drift on
machines with no AMD hardware.

Known-broken windows-msvc-* are the PR-only arm (#584), not this change.

Carried forward: #523's per-call hipMalloc/hipFree and hipStreamSynchronize are
illegal under hipGraph capture and must be reconciled before this capability
could ever be switched on.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants